W2 (KERNEL-SSM-MAMBA): the Mamba2 SSD CUDA arm — 9/9 mutations caught against the corrected bound (#496) - #675

…he gate host after a worktree loss (#496)

FOLLOWING_AGENTS_PROTOCOL

NOT AUTHORED BY THE COMMITTER. This is an operator recovery of a fresh
implementer's work after the fourth external deletion of an isolation worktree
this session. The implementer had built and run this green on both boxes and
staged it; the worktree was removed before the commit. The bytes survived on
the gate host and are restored here unchanged: md5 cbb1f928f4 for
cuda_mamba2_ssd.cuh and 37a0404433 for cuda_gdn.cu, matching what the
implementer reported before the loss.

The declared equivalence contract, which the implementer decided BEFORE writing
the kernel and recorded in the kernel header and all three test headers:

The CUDA arm keeps f32 accumulation throughout and deliberately does NOT mirror
upstream's tile downcasts. Those casts -- b.to(x_ptr.dtype.element_ty) at
ssd_chunk_state.py:283-285, cb.to(...)/prev_states.to(...) at
ssd_chunk_scan.py:266-269,359-363 -- are the input-precision requirement of
tl.dot, a tensor-core MMA. Every one of those tiles is loaded .to(tl.float32)
and computed in f32 right up to the MMA. These are scalar-FMA kernels with no
MMA, so mirroring the downcast would copy a constraint we do not have. The
inter-chunk `passed` buffer is allocated at state_dtype, NOT the host arm's f32
working width that spec 8.2 F9 warned W2 must not inherit.

A byte compare against the host arm is NOT reachable, and the downcasts are not
why: the two arms call different libms (CUDA expf <= 2 ulp, glibc <= 0.5).
Everything else is identical by construction. So the primary gate is the device
output against the same double-precision sequential reference at the same
upstream-ported tolerances the host arm uses, on the same inputs -- which
separates "device defect" from "wrong threshold". The derived device-vs-host
bar is rtol(K) = 4*(K+2)*2^-24, derived from 2.5 ulp of libm disagreement per
decay factor through a product of at most K plus (K-1)*u summation error. No
number was tuned and no tolerance was widened; each comparison logs the
fraction of budget actually used through MESSAGE rather than INFO, because
doctest prints INFO only on failure and an unaudited bar would have been a
false claim.

Evidence already captured on the gate host: Release build for 121a with CUTLASS
4.5.0, fa2 ENABLED and Marlin NVFP4 enabled, 0 warnings; RED run SIGSEGV on all
three binaries; GREEN run ssd 11 cases / 2069 assertions, state_update 10 /
5965, gated_norm 12 / 3723, all Status SUCCESS, exit 0, with zero
reference-tier lines.

That RED SIGSEGV is a real shared-seam defect, filed as #547 and deliberately
not fixed in flow: GB10 reports Backend::UnifiedMemory() == true, so
ReferenceTierEligible(kCUDA) is true, and with no native kernel GetOp installs
the CPU host kernel as a vt-cpu-ref provider and runs it over cudaMalloc
pointers. include/vt/backend.h already says a cudaMalloc pointer is not
host-dereferenceable on GB10; op_provider.cpp:515-526 gates on UnifiedMemory()
where it needs DeviceMemoryIsHostAddressable(). Every CUDA case here now calls
RequireNativeCudaProvider, so a device arm can never be gated by running the
host arm twice.

STILL OWED, and this branch is NOT landable until a fresh implementer finishes
it: the 8 scripted CUDA mutations, compute-sanitizer, the Debug arm, a full
ctest on the gate host, the spec section 8.3 that records the contract above
(its only copy was a staged blob and is presumed lost), and an origin/main
re-merge. A fresh review follows that, not this commit.

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]

…e $8.2 decode SUBCASE (#496)

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]

…aught, sanitizer clean, Debug arm green (#496)

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]

…tributed failure (#496)

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]

… -- a second job locks /tmp/gpu.lock (#496)

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]

…e attribution re-run is REMOTE_UNVERIFIED (#496)

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]

…, and the memory-safety claim outran the kernels (#496)

Tightening pass on PR #566 (`row/KERNEL-SSM-MAMBA-SSD-W2-FINISH` @ `1e819144e`),
whose fresh review returned PASS with five findings. The review's own work --
the equivalence contract's MMA reasoning, the reformulated mutations on CPU
twins, the grep, the recovered-byte md5s, the CI baseline subtraction -- is not
re-done here.

F1 (MEDIUM) -- the derivation, not the number, was wrong.

`cuda_mamba2_ssd.cuh` and the test headers both claimed the elementary functions
were the ONLY admitted source of device-vs-host divergence. They are not. Host
C++ is pinned `-ffp-contract=off` (CMakeLists.txt:55) so `a*b + c` keeps two
roundings; nothing passes `--fmad=false` to nvcc, so the header compiles at the
default `--fmad=true` and every `acc += a*b` in it is a single-rounding `fma`
whose host twin is not. This repo has MEASURED that idiom already
(.agents/benchmark-record.md:532). CMakeLists.txt:41-56 carves CUDA out of the
contraction policy because "GPU parity tests compare GPU-vs-GPU"; G2 is exactly
the case that carve-out does not cover.

The arithmetic, which the old constant did not survive:

  libm         2.5*K*u   (<= 2.5 ulp per decay factor, CUDA expf <= 2, glibc <= 0.5)
  summation    (K-1)*u   (length-K f32 sum; this is what amplifies the libm term)
  contraction  K*u       (the K product roundings the host keeps and fma does not)
  ----------------------------------------------------------------------------
  total        4.5*K*u - u

  old  4*(K+2)*u:  4.5K - 1 <= 4K + 8  <=>  K <= 18.  NOT PROVABLE at the
                   driver shapes, which run at K = T = 200.
  new  5*(K+2)*u:  4.5K - 1 <= 5K + 10 <=>  0.5K + 11 >= 0.  All K >= 0.

`-fmad=false` was weighed and REJECTED, not overlooked: nvcc takes it per
translation unit and this is a header included by `cuda_gdn.cu:48`, so applying
it means de-contracting every GDN decode kernel in that TU -- a measured hot
path -- or splitting a new `src/vt/` TU, which §8.3 already records as blocked on
#515. Slowing a shipped kernel to make a bound's prose true is the wrong trade.

Nothing was hidden numerically: re-scaling §8.4's audit by 4/5, the worst of 55
comparisons goes 7.66% -> 6.13% of budget, the driver shapes 0.32%/0.18% ->
0.26%/0.14%, and mutant M3 962173% -> 769738%, still caught by four orders of
magnitude. The bound moved because the DERIVATION gained a term the build
actually emits, and the header comment, all three test comments and spec §8.3
now say the same true thing.

F2 (MEDIUM) -- both halves taken: the free clamps AND the narrowed claim.

The shared validator checks metadata shape/dtype/device only (`CheckI32Meta`,
ops.cpp:1717-1723); every VALUE check lives in the CPU kernel (cpu_ops.cpp:1622-
1648), so the device arm silently drops six of them. Three are memory-unsafe,
and the stated reason for dropping them -- a D2H plus a stream sync -- does not
apply, because the values are already in device registers and the decode kernel
has always clamped its `state_indices` slot for free on that basis.

  * `M2StatePassKernel`: `chunk_end` clamped to `nchunks`, `chunk_start` to 0.
    Unclamped, `lci[b] >= nchunks` makes `M2Store(passed, ...)` an out-of-bounds
    WRITE past the `cudaMallocAsync` allocation (W1 finding F7's device half).
  * `M2ChunkScanKernel`: `si_ok = si >= 0 && si < S`, and `!si_ok` opens the
    chunk from a zero previous state. Unclamped, an out-of-range `seq_idx[c]`
    reads `initial_states` out of bounds, and a `seq_idx[0] < 0` additionally
    makes `si == si_prev` at c == 0 and indexes `passed` at chunk -1 -- a hole
    the finding did not name and this pass found while writing the clamp. The
    kernel gained an `S` parameter for it.

In-contract behaviour is bit-identical: in contract `si` is always in range and
`lci` always below `nchunks`, so neither clamp can fire.

The claim is narrowed at the same time, because clamping two does not make the
arm memory-safe. The `cu_chunk_seqlens` tiling and per-chunk length checks are
NOT clamped and a violation IS memory-unsafe -- a garbage `ccs` indexes
x/B/C/z/out out of bounds in every stage -- and the header and §8.3 now say so
instead of folding it into a blanket "the device kernels remain MEMORY SAFE".

Pinned, not asserted in prose: a device-only case runs both violations against
in-contract reference runs whose result each clamp is DEFINED to reproduce, so
the assertions are exact rather than tolerances.

F3, F4 (LOW, record accuracy). M7's §8.4 label overstated the device mutant: the
launcher's `nblocks = rows * args.n_groups` is not mutated while the kernel's
`n_groups` is forced to 1, so blocks `blk >= rows` run past the tensor and the
mutant is memory-unsafe, failing partly for that rather than purely on whole-row
variance. The guarantee IS pinned by the reviewer's clean CPU twin; only the
label was wrong. §8.2's residual sentence still gave the DOWNCASTS as the reason
W2 cannot byte-compare, written when W2 was expected to mirror them; §8.3
supersedes it -- both arms stay f32 and the reasons are libm and contraction.
Two sentences in one spec gave two causes for one fact; reconciled.

F5 (LOW). Taken. The five per-call scratch buffers are held by an `M2Scratch`
scope guard, so a throw on the Nth `cudaMallocAsync` no longer leaks the N-1
before it, and `Release()` frees all five before reporting rather than leaking
the remainder on a mid-sequence free failure.

EVIDENCE. `df -h /` 85% used / 67G free before and after every result.

  test_ops_mamba2_ssd            8/8    1175/1175  SUCCESS!
  test_ops_mamba2_state_update   6/6    2469/2469  SUCCESS!
  test_ops_mamba2_gated_norm     9/9    2107/2107  SUCCESS!

Identical to the pre-change counts, as expected -- every code change is inside
`#ifdef VLLM_CPP_CUDA` or in the `.cuh`. `Status:` was read, not `assertions:`
alone.

OMITTED_GATES -- the CUDA arm was neither built nor run. `dgx.casa` has been
unreachable since 06:50 CEST and this box has no nvcc and no GPU. Two
substitutes were run and neither is offered as the device gate:

  1. The `.cuh` compiled at `-std=c++20 -Wall -Wextra -Werror` against CUDA
     shims with each `Kernel<<<cfg>>>(args)` rewritten to `M2Sink(cfg),
     Kernel(args)` -- dropping the launch config while PRESERVING the arity and
     type check on all 7 launches. EXIT 0. Proved ARMED by deleting the `S`
     argument on a scratch copy: `too few arguments to function
     M2ChunkScanKernel`, exit 1.
  2. A CPU twin of the two clamped index computations, over the device case's
     own shape. Unclamped, every claimed hole reproduced: index 2432 and -4096
     into a 2048-element `passed`, 463232 into a 2048-element `initial_states`,
     -512 for the c == 0 hole. Clamped, all land in [0,1920], the `lci` clamp
     reproduces the in-contract index range exactly, and both `seq_idx`
     violations read no previous state at all.

Still owed on device: the three CUDA arms, `compute-sanitizer memcheck` on the
new case (which is what actually proves memory safety -- an out-of-bounds write
into a pool allocation commonly does not fault), a mutation re-sweep against the
moved bound, and §8.4's `~/w2ssd/refail.log`, still REMOTE_UNVERIFIED.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]

…g pass (#496)

Clean Release rebuild of every target -- 817/817 ninja edges, 0 warnings, 0
errors at `-Wall -Wextra -Werror` -- then full `ctest -j 4`: 100% tests passed,
0 failed out of 403, CTEST_EXIT=0, 22.35 s, 2 skipped. `df -h /` 87% used /
60G free.

Recorded with the caveat that matters: this is the CPU-ONLY lane, so it is a
much smaller and faster gate than §8.4's 431-test GPU-host run and is not a
substitute for it. None of §8.4's ten failures is reachable from a build with
no CUDA. The device arm remains `omitted_gates` while `dgx.casa` is
unreachable.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]

…D bound (#496)

§8.4 reports "the worst one used 7.66% of rtol(K) = 4*(K+2)*2^-24". That is
still what the run produced, but §8.5 moved the bound to 5*(K+2)*2^-24, so read
without a pointer it now looks like a statement about the current bar.

Cross-referenced rather than restated: the captured numbers stay as captured,
with the conversion (6.13%, 0.26%/0.14%, 769738%) named next to them and the
arithmetic in §8.3 point 6. Rewriting a measurement to match a later derivation
is how a record stops being evidence.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]

…ed the gap it does not cover (#496)

Two inaccuracies in the comment written one commit ago, found re-reading it:

It said "two of the dropped ones are memory-UNSAFE", which reads as "and the
other four are not". The header enumerates six, of which the `cu_chunk_seqlens`
tiling and per-chunk length checks are ALSO memory-unsafe and are NOT clamped.
The two this case covers are the two that are memory-unsafe AND bounded by a
register-local clamp. Narrowing the count in the very comment that exists to
stop a claim outrunning the kernels would have re-introduced F2 at a smaller
scale, so the comment now points at the header's full list and names the
uncovered gap explicitly.

It also cited §8.4 for the owed `compute-sanitizer memcheck`; that is recorded
in §8.5.

Comment-only. Rebuilt clean (392/392, 0 warnings) and re-ran: 8/8, 1175/1175,
`Status: SUCCESS!`. The `#ifdef VLLM_CPP_CUDA` region was re-checked with a
`-DVLLM_CPP_CUDA -fsyntax-only` compile at -Wall -Wextra -Werror, exit 0.
`df -h /` 83% used / 73G free.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]

… reproduced standalone on an idle box (#496)

FOLLOWING_AGENTS_PROTOCOL

The last unattributed ctest failure is closed, and the contention hypothesis
this spec had recorded as "plausible, not proven" is REFUTED rather than quietly
kept: test_minimax_h3 reproduces STANDALONE, serially, under the lock, on a box
that had just rebooted and had no CUDA process resident. It is #486 (open) with
#516 as its root cause -- which is what row/pool-device-key is repairing, and
therefore why that baseline branch passed a test main-based branches fail. The
difference was the baseline carrying a FIX, not this branch carrying a defect.
The re-run reproduces #486's recorded signature number for number:
38 | 36 passed | 2 failed | 41 skipped, 42724 assertions, the same two case
names, the same `cudaFree: invalid argument`, the same SIGSEGV.

All ten ctest failures are now pre-existing and tracked. None is W2's.

Also fixed a live protocol defect found while doing it: the re-run had been
relaunched as `flock -w 3600 $HOME/gpu.lock ./w2refail.sh`, an OUTER flock
wrapping a script that takes the SAME lock on its own fd. flock locks an open
file DESCRIPTION, so the inner acquisition blocked against its own parent -- a
self-deadlock HOLDING the shared GPU mutex with three other agents' jobs queued
behind it. Killed the stack, the lock passed straight to a waiter, and relaunched
with a single acquisition.

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]

…- deterministic, not starved (#496)

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]

…vidence; 32G build tree reclaimed (#496)

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]

…erator-run on dgx, not CI (#496)

Two corrections to §8.5, both narrowing what this branch claims.

1. THE CUDA COMPILE AXIS IS CLOSED, AND IT WAS NOT MY RESULT.

`dgx.casa` returned at 06:57 UTC after a reboot, which supersedes §8.4's
REMOTE_UNVERIFIED note on the host being away. The operator transferred this
branch by `git archive` -- never rsync, which has previously overwritten
goldens into false passes -- and built it with real nvcc:

  NVCC: cuda_13.0.r13.0
  CUTLASS found at ~/cutlass-4.5.0; enabling sm120a NVFP4 cutlass GEMM
  Marlin NVFP4 W4A16 MoE GEMM enabled (vendored) for [121a]
  FlashAttention-2 prefill/decode: ENABLED for arch(es) [121a]
  CONFIGURE_EXIT=0  BUILD_EXIT=0  WARNINGS=0  ENOSPC=0

Zero errors, zero warnings under the project's -Werror flags, disk unchanged at
64G either side so this is not the stale-binary false green, and the three
fast-path features READ OUT OF the configure log rather than assumed -- an
absent CUTLASS exits 0 too, so "the build succeeded" alone proves nothing.

That retires the risk my shim arity-check could only approximate: the new `S`
kernel parameter and the `M2Scratch` guard compile through `cuda_gdn.cu` at the
arch this ships on. Recorded as OPERATOR-RUN. I did not run it, this box has no
nvcc, and `cuda-fat-build` has still never completed on this branch -- so no CI
job has compiled this code either. The shim check and the F2 clamp CPU twin are
kept as what an implementer could establish unaided, not restated as the gate.

2. THE REMAINING GATES ARE BLOCKED BY THE GPU LOCK, NOT BY THE HOST.

The old wording said "omitted_gates until dgx.casa returns". It has returned,
and they are still owed -- `$HOME/gpu.lock` is held by other coordinators' jobs
with 8h timeouts. A reachable host is not an available GPU, and running these
against a contended one reproduces exactly the undetected-contention defect
§8.4 already records. Owed: the three CUDA arms (EXECUTION, which a compile
does not supply); `compute-sanitizer memcheck` on the new clamp case, which is
what actually proves F2 because an out-of-bounds write into a cudaMallocAsync
pool commonly does not fault; the 9-mutation re-sweep against the MOVED bound,
since a widened bound is precisely the change that could stop a mutation
reddening; and §8.4's refail.log.

Item 3 is the one a reader is most likely to wave through. The re-scaled margins
in §8.3 point 6 make it very likely to hold, and very likely is not a result.

3. THE CANCELLED-vs-FAILED TRAP, WHICH TWO OF US HIT FROM OPPOSITE ENDS.

Every Actions run for #592 -- four SHAs, both workflows -- ended
`conclusion: cancelled`. Three were my own follow-up pushes. The fourth was
killed at 07:44:45-07:44:55 together with EVERY run in the repository, 20 of 20
across 7 branches, `windows-msvc-cpu` dying mid-build after passing two steps.
That is an Actions-side event, not a verdict on any diff.

`gh pr checks` renders a cancelled job as `fail`. My watcher reported 16
failures and the operator's reported 20; the true count both times was ZERO.
A per-check listing cannot separate "this branch is red" from "the pool was
killed" -- only the run-level `conclusion` can. Carried in the spec because
anyone subtracting a CI baseline on this repo will otherwise attribute an
infrastructure event to a diff.

A re-run was triggered at 07:47 and was still queued at 07:50. Whatever it
reports is the CI result; this commit does not claim one.

Record-only; no code change. `df -h /` 84% used / 73G free.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]

FOLLOWING_AGENTS_PROTOCOL

Neither W2 branch contained the other. W2-FINISH carries the three evidence
commits that resolve the merge precondition -- test_minimax_h3 ATTRIBUTED to
#486/#516, all ten ctest failures reproducing STANDALONE, and the evidence-log
location. W2-FIX carries the F1/F2 repair and the operator-run nvcc compile.
Both are needed; this consolidates them so one squash lands the row.

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]

FOLLOWING_AGENTS_PROTOCOL

Routine re-merge before landing. `git merge --no-edit` writes git's default
message with no trailers, which check-commit-trailers rejects -- it exempts
nothing, merges included. Amended rather than left, because the same
trailer-less merge commit on a sibling branch is exactly what reddens the
push lane when a branch is landed with a local --no-ff instead of a squash.

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]

…ement (#496)

FOLLOWING_AGENTS_PROTOCOL

Spec 8.6. The four items 8.5 listed as owed are closed -- three by measurement
on the gate host, one by attribution.

The CUDA arms and compute-sanitizer ran under one lock acquisition: ssd
12/2095, state_update 10/5965, gated_norm 12/3723, all SUCCESS, memcheck
ERROR SUMMARY 0 errors on both. The assertion counts are quoted because a suite
that skipped every CUDA case would also exit 0; a non-zero count is what
separates "the device arms ran" from "the binary started".

The re-sweep ran against the MOVED bound, 5*(K+2)*u, which is the whole point:
8.4's sweep was scored against 4*(K+2)*u and a widened bound is precisely the
change that could stop a mutation reddening. Controls first at the same counts
as the arms, so an empty run is excluded. 9 of 9 CAUGHT. Every mutant built in
a phase that takes no lock, run under one acquisition, source restored and
md5-verified after each. Every form keeps every variable read, because a
mutation that fails to compile under -Werror=all-warnings is not a caught
mutation, it is a suite that never ran -- two of the original eight were exactly
that.

M6 is the one worth reading: it aborts while printing
"assertions: 2577 | 2577 passed | 0 failed", a clean assertions line on a
FAILING run, and is caught only because the harness reads the exit code. That is
this repo's recurring trap in its sharpest form, and it is why item 3 could not
have been discharged by inspecting summaries.

test_minimax_h3 is ATTRIBUTED, not waived: reproduced standalone on an idle box
under the lock, TEST_EXIT=139, and it is #486 with root cause #516,
signature-for-signature. The independent baseline that PASSED was
row/pool-device-key, the branch that FIXES #516 -- so the baseline carried a fix,
this branch does not carry a defect. The contention hypothesis 8.4 recorded as
"plausible, not proven" is REFUTED rather than quietly retained.

Records only, no product code touched.

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]

FOLLOWING_AGENTS_PROTOCOL

main moved twice more during the gate run. Merged with an explicit message
because `git merge --no-edit` writes git's default, which check-commit-trailers
rejects -- it exempts nothing, merges included.

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]